


void f_FitWeight(
                 __global float* nodeVarA,
                 __global float* nodeVarB,
                 __global float16* WeightsNode,
                 __global uint2*   WeightOffset_size,   //权重偏移 和输入节点数量
                 
                 __global float*    nodeVar_Buf,         //调整后的 节点值
                 __global uchar*    WeightModifyID,        //权重偏移 和输入节点数量

                 float maxDeviation, float LearningRate, ushort BlockNum, ushort layer
);




__kernel void f_aiHiddenLayerFeedforward(
                                        __global float*   nodeVarA,
                                        __global float*   nodeVarB,
                                        __global float16* WeightsNode,
                                        __global uint2*   WeightOffset_size,
                                        
                                        __global float*   nodeVar_Buf,
                                        __global uchar*   WeightModifyID,        //权重偏移 和输入节点数量

                                        float   maxDeviation, float LearningRate, ushort BlockNum, ushort layer

) {

    size_t groupID = get_global_linear_id();
    //if(groupID >= WeightOffset_size[layer].s1) return;
    
    size_t Out_nodeVar_ID = groupID + WeightOffset_size[layer].s0;
    nodeVarB[groupID] = nodeVar_Buf[Out_nodeVar_ID];
    //printf("nodeVarB[groupID] %v2f \n", (float2)(nodeVarB[groupID], Out_nodeVar_ID));
    //printf("Out_nodeVar_ID %f \n", nodeVarB[0]);
    
    if(groupID==0){
        size_t localID = get_local_linear_id();
        //size_t In_nodeVar_ID = localID + WeightOffset_size[layer-1].s0;
        /* ----------------- 层A 输入层数据修正计算---------------------
        层A 层W 层O
         . .    .
           .

         . .    .
           .

         . .
           .

           修正数据拷贝到 nodeVar_Buf and nodeVarA
        * ---------------------------------------begin----------------------------------------*/
        //float16 tmpDeviation;   //当前层 将要计算的上一层相对于当前输入层误差
        float16 a = (float16)(1,1,1,1,1,1,1,1,1,1,1,1,1,1,1,1);
        ushort LocalBlockNum = WeightOffset_size[layer-1].s1/16;
        for(uchar i=0; i<LocalBlockNum; ++i){
            size_t In_nodeVar_ID = 0;
            if(layer){
                //每隔权重索引的16个偏移 采用ID
                In_nodeVar_ID = i + WeightOffset_size[layer-1].s0 / (16*BlockNum);

                float16 nodeVar = vload16(In_nodeVar_ID, nodeVar_Buf);

                //printf("WeightOffset_size ----------- %u \n", i);
                //printf("new_outNodeVar %v16f \n", nodeVar);
                //计算当前输入层新值
                float16 tmp_In_NodeNewVar = nodeVar * (maxDeviation);


                /*uint modify_wid = WeightModifyID[WeightOffset_size[layer].s1 + groupID];
                switch( modify_wid ){
                    case 0:
                        nodeVar.s0 *= maxDeviation; //if(tmpWeigth[i].s0 > 1.0f) tmpWeigth[i].s0 = 1.0f; else if(tmpWeigth[i].s0 < -1.0f) tmpWeigth[i].s0 = -1.0f;
                        break;
                    case 1:
                        nodeVar.s1 *= maxDeviation; //if(tmpWeigth[i].s1 > 1.0f) tmpWeigth[i].s1 = 1.0f; else if(tmpWeigth[i].s1 < -1.0f) tmpWeigth[i].s1 = -1.0f;
                        //printf("tmpWeigth %i \n", modify_wid);
                        break;
                    case 2:
                        nodeVar.s2 *= maxDeviation; //if(tmpWeigth[i].s2 > 1.0f) tmpWeigth[i].s2 = 1.0f; else if(tmpWeigth[i].s2 < -1.0f) tmpWeigth[i].s2 = -1.0f;
                        //printf("tmpWeigth %i \n", modify_wid);
                        break;
                    case 3:
                        nodeVar.s3 *= maxDeviation; //if(tmpWeigth[i].s3 > 1.0f) tmpWeigth[i].s3 = 1.0f; else if(tmpWeigth[i].s3 < -1.0f) tmpWeigth[i].s3 = -1.0f;
                        //printf("tmpWeigth %i \n", modify_wid);
                        break;
                    case 4:
                        nodeVar.s4 *= maxDeviation; //if(tmpWeigth[i].s4 > 1.0f) tmpWeigth[i].s4 = 1.0f; else if(tmpWeigth[i].s4 < -1.0f) tmpWeigth[i].s4 = -1.0f;
                        //printf("tmpWeigth %i \n", modify_wid);
                        break;
                    case 5:
                        nodeVar.s5 *= maxDeviation; //if(tmpWeigth[i].s5 > 1.0f) tmpWeigth[i].s5 = 1.0f; else if(tmpWeigth[i].s5 < -1.0f) tmpWeigth[i].s5 = -1.0f;
                        //printf("tmpWeigth %i \n", modify_wid);
                        break;
                    case 6:
                        nodeVar.s6 *= maxDeviation; //if(tmpWeigth[i].s6 > 1.0f) tmpWeigth[i].s6 = 1.0f; else if(tmpWeigth[i].s6 < -1.0f) tmpWeigth[i].s6 = -1.0f;
                        //printf("tmpWeigth %i \n", modify_wid);
                        break;
                    case 7:
                        nodeVar.s7 *= maxDeviation; //if(tmpWeigth[i].s7 > 1.0f) tmpWeigth[i].s7 = 1.0f; else if(tmpWeigth[i].s7 < -1.0f) tmpWeigth[i].s7 = -1.0f;
                        //printf("tmpWeigth %i \n", modify_wid);
                        break;
                    case 8:
                        nodeVar.s8 *= maxDeviation; //if(tmpWeigth[i].s8 > 1.0f) tmpWeigth[i].s8 = 1.0f; else if(tmpWeigth[i].s8 < -1.0f) tmpWeigth[i].s8 = -1.0f;
                        //printf("tmpWeigth %i \n", modify_wid);
                        break;
                    case 9:
                        nodeVar.s9 *= maxDeviation; //if(tmpWeigth[i].s9 > 1.0f) tmpWeigth[i].s9 = 1.0f; else if(tmpWeigth[i].s9 < -1.0f) tmpWeigth[i].s9 = -1.0f;
                        //printf("tmpWeigth %i \n", modify_wid);
                        break;
                    case 10:
                        nodeVar.sa *= maxDeviation; //if(tmpWeigth[i].sa > 1.0f) tmpWeigth[i].sa = 1.0f; else if(tmpWeigth[i].sa < -1.0f) tmpWeigth[i].sa = -1.0f;
                        //printf("tmpWeigth %i \n", modify_wid);
                        break;
                    case 11:
                        nodeVar.sb *= maxDeviation; //if(tmpWeigth[i].sb > 1.0f) tmpWeigth[i].sb = 1.0f; else if(tmpWeigth[i].sb < -1.0f) tmpWeigth[i].sb = -1.0f;
                        //printf("tmpWeigth %i \n", modify_wid);
                        break;
                    case 12:
                        nodeVar.sc *= maxDeviation; //if(tmpWeigth[i].sc > 1.0f) tmpWeigth[i].sc = 1.0f; else if(tmpWeigth[i].sc < -1.0f) tmpWeigth[i].sc = -1.0f;
                        //printf("tmpWeigth %i \n", modify_wid);
                        break;
                    case 13:
                        nodeVar.sd *= maxDeviation; //if(tmpWeigth[i].sd > 1.0f) tmpWeigth[i].sd = 1.0f; else if(tmpWeigth[i].sd < -1.0f) tmpWeigth[i].sd = -1.0f;
                        //printf("tmpWeigth %i \n", modify_wid);
                        break;
                    case 14:
                        nodeVar.se *= maxDeviation; //if(tmpWeigth[i].se > 1.0f) tmpWeigth[i].se = 1.0f; else if(tmpWeigth[i].se < -1.0f) tmpWeigth[i].se = -1.0f;
                        //printf("tmpWeigth %i \n", modify_wid);
                        break;
                    case 15:
                        nodeVar.sf *= maxDeviation; //if(tmpWeigth[i].sf > 1.0f) tmpWeigth[i].sf = 1.0f; else if(tmpWeigth[i].sf < -1.0f) tmpWeigth[i].sf = -1.0f;
                        //printf("tmpWeigth %i \n", modify_wid);
                        break;
                }
                
                ++WeightModifyID[WeightOffset_size[layer].s1 + groupID];*/
                
                
                //nodeVar *= maxDeviation;
               /* nodeVar.s0 = 1.0 / (nodeVar.s0+0.000000000001);
                nodeVar.s1 = 1.0 / (nodeVar.s1+0.000000000001);
                nodeVar.s2 = 1.0 / (nodeVar.s2+0.000000000001);
                nodeVar.s3 = 1.0 / (nodeVar.s3+0.000000000001);
                nodeVar.s4 = 1.0 / (nodeVar.s4+0.000000000001);
                nodeVar.s5 = 1.0 / (nodeVar.s5+0.000000000001);
                nodeVar.s6 = 1.0 / (nodeVar.s6+0.000000000001);
                nodeVar.s7 = 1.0 / (nodeVar.s7+0.000000000001);
                nodeVar.s8 = 1.0 / (nodeVar.s8+0.000000000001);
                nodeVar.s9 = 1.0 / (nodeVar.s9+0.000000000001);
                nodeVar.sa = 1.0 / (nodeVar.sa+0.000000000001);
                nodeVar.sb = 1.0 / (nodeVar.sb+0.000000000001);
                nodeVar.sc = 1.0 / (nodeVar.sc+0.000000000001);
                nodeVar.sd = 1.0 / (nodeVar.sd+0.000000000001);
                nodeVar.se = 1.0 / (nodeVar.se+0.000000000001);
                nodeVar.sf = 1.0 / (nodeVar.sf+0.000000000001);*/
                
                //pow(nodeVar, );

                //barrier(CLK_LOCAL_MEM_FENCE);
                //printf("--- nodeVar %v16f \n", nodeVar);
                //printf("A -- new_outNodeVar %2.2v16f \n", tmp_In_NodeNewVar);
                //printf("A -------------  maxDeviation %i \n", BlockNum);
                //printf("A -------------  LocalBlockNum = %u \n", LocalBlockNum);
                //printf("A -------------  maxDeviation = %f \n", maxDeviation);
                

                //误差值累加
                //a += tmpDeviation;
                //新输入值拷贝到合并缓存中
                vstore16(nodeVar, In_nodeVar_ID, nodeVar_Buf);
                vstore16(nodeVar, i, nodeVarA);


                //计算当前输入层误差
                nodeVar.s0 += 1.0;
                nodeVar.s1 += 1.0;
                nodeVar.s2 += 1.0;
                nodeVar.s3 += 1.0;
                nodeVar.s4 += 1.0;
                nodeVar.s5 += 1.0;
                nodeVar.s6 += 1.0;
                nodeVar.s7 += 1.0;
                nodeVar.s8 += 1.0;
                nodeVar.s9 += 1.0;
                nodeVar.sa += 1.0;
                nodeVar.sb += 1.0;
                nodeVar.sc += 1.0;
                nodeVar.sd += 1.0;
                nodeVar.se += 1.0;
                nodeVar.sf += 1.0;

                tmp_In_NodeNewVar.s0 += 1.0;
                tmp_In_NodeNewVar.s1 += 1.0;
                tmp_In_NodeNewVar.s2 += 1.0;
                tmp_In_NodeNewVar.s3 += 1.0;
                tmp_In_NodeNewVar.s4 += 1.0;
                tmp_In_NodeNewVar.s5 += 1.0;
                tmp_In_NodeNewVar.s6 += 1.0;
                tmp_In_NodeNewVar.s7 += 1.0;
                tmp_In_NodeNewVar.s8 += 1.0;
                tmp_In_NodeNewVar.s9 += 1.0;
                tmp_In_NodeNewVar.sa += 1.0;
                tmp_In_NodeNewVar.sb += 1.0;
                tmp_In_NodeNewVar.sc += 1.0;
                tmp_In_NodeNewVar.sd += 1.0;
                tmp_In_NodeNewVar.se += 1.0;
                tmp_In_NodeNewVar.sf += 1.0;


                //a =  nodeVar / tmp_In_NodeNewVar;
                a =  tmp_In_NodeNewVar / nodeVar;

                
            }else{
                In_nodeVar_ID = i;
                float16 nodeVar = vload16(In_nodeVar_ID, nodeVar_Buf);
                vstore16(nodeVar, i, nodeVarA);
            }
            
            
            //float t1[32]; for(uint i=0; i<32; ++i) t1[i] = i;
            //printf("- ar %v16f \n", tmp_In_NodeNewVar);
            //for(uint i=0; i<16; ++i) printf("n-A %f  ", nodeVarA[i]);
            //printf("n \n");
            //for(uint i=0; i<16; ++i) printf("nodeVarA %f \n", nodeVarA[ i ]);
            //barrier(CLK_GLOBAL_MEM_FENCE);
        }
        //---------------------------------------end--------------------------------------------
        
        /*float t1[32];
        for(uint i=0; i<32; ++i) t1[i] = i;
        float16 t2 = (float16)(77,77,77,77,77,77,77,77,77,77,77,77,77,77,77,77);
        float16 t_16 = vload16(1, t1);
        printf("------------- t_16 ---------------- %v16f \n", t_16);*/

        //-------------------计算输入层 最大误差-----------------------
        /*float2 b = (float2)(a.s0, a.s0);
        b.s0 = max(b.s0, a.s1);
        b.s0 = max(b.s0, a.s2);
        b.s0 = max(b.s0, a.s3);
        b.s0 = max(b.s0, a.s4);
        b.s0 = max(b.s0, a.s5);
        b.s0 = max(b.s0, a.s6);
        b.s0 = max(b.s0, a.s7);
        b.s0 = max(b.s0, a.s8);
        b.s0 = max(b.s0, a.s9);
        b.s0 = max(b.s0, a.sa);
        b.s0 = max(b.s0, a.sb);
        b.s0 = max(b.s0, a.sc);
        b.s0 = max(b.s0, a.sd);
        b.s0 = max(b.s0, a.se);
        b.s0 = max(b.s0, a.sf);

        b.s1 = min(b.s1, a.s1);
        b.s1 = min(b.s1, a.s2);
        b.s1 = min(b.s1, a.s3);
        b.s1 = min(b.s1, a.s4);
        b.s1 = min(b.s1, a.s5);
        b.s1 = min(b.s1, a.s6);
        b.s1 = min(b.s1, a.s7);
        b.s1 = min(b.s1, a.s8);
        b.s1 = min(b.s1, a.s9);
        b.s1 = min(b.s1, a.sa);
        b.s1 = min(b.s1, a.sb);
        b.s1 = min(b.s1, a.sc);
        b.s1 = min(b.s1, a.sd);
        b.s1 = min(b.s1, a.se);
        b.s1 = min(b.s1, a.sf);*/
        //float newMaxDeviation = (b.s0 + b.s1);
        float newMaxDeviation = 1.0 / ((a.s0+a.s1+a.s2+a.s3+a.s4+a.s5+a.s6+a.s7+a.s8+a.s9+a.sa+a.sb+a.sc+a.sd+a.se+a.sf));
        
        //误差为0 则不调节上一层权重
        if(groupID == 0 && newMaxDeviation != 0.0){
            //newMaxDeviation += sign(newMaxDeviation);
            //if(fabs(newMaxDeviation) < 0.0000000000000001) newMaxDeviation = 0.0;
            
            ushort LocalBlockNum = WeightOffset_size[layer-1].s1/16;
            if(!LocalBlockNum) LocalBlockNum = 1;

            ushort GroupNum_OutLayer = max((ushort)(WeightOffset_size[layer].s1), LocalBlockNum);

            //ushort2 range = (ushort2)(GroupNum_OutLayer, LocalBlockNum);
            //printf("B -------------------  new_outNodeVar %f \n", newMaxDeviation);
            //printf("Local = %u \n", LocalBlockNum);
            //printf("goud = %u \n", GroupNum_OutLayer);
            ndrange_t ndrange1 = ndrange_1D(GroupNum_OutLayer, LocalBlockNum);
            
            barrier(CLK_LOCAL_MEM_FENCE);
            barrier(CLK_GLOBAL_MEM_FENCE);

            void(^func_SamplingBlock)(void) = ^{ f_FitWeight(nodeVarA, nodeVarB, WeightsNode, WeightOffset_size, nodeVar_Buf, WeightModifyID, maxDeviation, LearningRate, LocalBlockNum, layer); } ;
            int err_ret = enqueue_kernel(get_default_queue(), CLK_ENQUEUE_FLAGS_WAIT_KERNEL, ndrange1, func_SamplingBlock);

            
            //printf("B -------------------  err_ret %u \n", err_ret);
        }
        
    }
    

}


//拟合权重
void f_FitWeight(
                 __global float* nodeVarA,
                 __global float* nodeVarB,
                 __global float16* WeightsNode,
                 __global uint2*   WeightOffset_size,   //权重偏移 和输入节点数量
                 
                 __global float*    nodeVar_Buf,         //调整后的 节点值
                 __global uchar*    WeightModifyID,        //权重偏移 和输入节点数量

                 float maxDeviation, float LearningRate, ushort BlockNum, ushort layer
) {
    size_t groupID = get_global_linear_id();
    if(groupID >= WeightOffset_size[layer].s1) return;

    size_t localID = get_local_linear_id();
    
    //取权重ID
    size_t weightsID = groupID*BlockNum + WeightOffset_size[layer].s0 + localID;
    //size_t weightsID = (WeightOffset_size[layer].s0) + localID;
    //printf("BlockNum %u \n", BlockNum);
    
    __local float16 tmpWeigth[32];
    tmpWeigth[localID] = WeightsNode[weightsID];
    //barrier(CLK_LOCAL_MEM_FENCE);
    //float16 tmpWeigth[16];
    //tmpWeigth[localID] = WeightsNode[weightsID];
    /*if(groupID==0) {
        if(localID == 0){
            printf("maxDeviation = %f \n", maxDeviation);
        }
        for(uchar i=0; i<BlockNum; ++i) {
            const float16 InNodeVar = vload16(i, nodeVarA);
            printf("B -- new_outNodeVar %2.4v16f \n", InNodeVar);
        }
    }*/
    
    if(localID == 0){
        uchar i=0;
        uint modify_wid = WeightModifyID[WeightOffset_size[layer].s1 + groupID];
        
        float outNodeVar = nodeVarB[groupID] * (1.0f-maxDeviation);
        
        
        float AdjusRatio = 0.0;
        for(uchar correctNum=0; correctNum<16; ++correctNum){
            float16 a;
            
            if(modify_wid >= 16) modify_wid = 0;
            for(i=0; i<BlockNum; ++i){
                const float16 InNodeVar = vload16(i, nodeVarA);
                
                //printf("tmpWeigth %#v16f \n", tmpWeigth[ i ]);
                //printf("modify_wid %u \n", modify_wid);
                switch( modify_wid ){
                    case 0:
                        tmpWeigth[i].s0 += AdjusRatio; //if(tmpWeigth[i].s0 > 1.0f) tmpWeigth[i].s0 = 1.0f; else if(tmpWeigth[i].s0 < -1.0f) tmpWeigth[i].s0 = -1.0f;
                        //tmpWeigth[i].s0 *= maxDeviation;
                        //printf("tmpWeigth %i \n", modify_wid);
                        break;
                    case 1:
                        
                        tmpWeigth[i].s1 += AdjusRatio; //if(tmpWeigth[i].s1 > 1.0f) tmpWeigth[i].s1 = 1.0f; else if(tmpWeigth[i].s1 < -1.0f) tmpWeigth[i].s1 = -1.0f;
                        //tmpWeigth[i].s1 *= maxDeviation;
                        //printf("tmpWeigth %i \n", modify_wid);
                        break;
                    case 2:
                        
                        tmpWeigth[i].s2 += AdjusRatio; //if(tmpWeigth[i].s2 > 1.0f) tmpWeigth[i].s2 = 1.0f; else if(tmpWeigth[i].s2 < -1.0f) tmpWeigth[i].s2 = -1.0f;
                        //tmpWeigth[i].s2 *= maxDeviation;
                        break;
                    case 3:
                        
                        tmpWeigth[i].s3 += AdjusRatio; //if(tmpWeigth[i].s3 > 1.0f) tmpWeigth[i].s3 = 1.0f; else if(tmpWeigth[i].s3 < -1.0f) tmpWeigth[i].s3 = -1.0f;
                        //tmpWeigth[i].s3 *= maxDeviation;
                        break;
                    case 4:
                        
                        tmpWeigth[i].s4 += AdjusRatio; //if(tmpWeigth[i].s4 > 1.0f) tmpWeigth[i].s4 = 1.0f; else if(tmpWeigth[i].s4 < -1.0f) tmpWeigth[i].s4 = -1.0f;
                        //tmpWeigth[i].s4 *= maxDeviation;
                        break;
                    case 5:
                        
                        tmpWeigth[i].s5 += AdjusRatio; //if(tmpWeigth[i].s5 > 1.0f) tmpWeigth[i].s5 = 1.0f; else if(tmpWeigth[i].s5 < -1.0f) tmpWeigth[i].s5 = -1.0f;
                        //tmpWeigth[i].s5 *= maxDeviation;
                        break;
                    case 6:
                        
                        tmpWeigth[i].s6 += AdjusRatio; //if(tmpWeigth[i].s6 > 1.0f) tmpWeigth[i].s6 = 1.0f; else if(tmpWeigth[i].s6 < -1.0f) tmpWeigth[i].s6 = -1.0f;
                        //tmpWeigth[i].s6 *= maxDeviation;
                        break;
                    case 7:
                        
                        tmpWeigth[i].s7 += AdjusRatio; //if(tmpWeigth[i].s7 > 1.0f) tmpWeigth[i].s7 = 1.0f; else if(tmpWeigth[i].s7 < -1.0f) tmpWeigth[i].s7 = -1.0f;
                        //tmpWeigth[i].s7 *= maxDeviation;
                        break;
                    case 8:
                        
                        tmpWeigth[i].s8 += AdjusRatio; //if(tmpWeigth[i].s8 > 1.0f) tmpWeigth[i].s8 = 1.0f; else if(tmpWeigth[i].s8 < -1.0f) tmpWeigth[i].s8 = -1.0f;
                        //tmpWeigth[i].s8 *= maxDeviation;
                        break;
                    case 9:
                        
                        tmpWeigth[i].s9 += AdjusRatio; //if(tmpWeigth[i].s9 > 1.0f) tmpWeigth[i].s9 = 1.0f; else if(tmpWeigth[i].s9 < -1.0f) tmpWeigth[i].s9 = -1.0f;
                        //tmpWeigth[i].s0 *= maxDeviation;
                        break;
                    case 10:
                        
                        tmpWeigth[i].sa += AdjusRatio; //if(tmpWeigth[i].sa > 1.0f) tmpWeigth[i].sa = 1.0f; else if(tmpWeigth[i].sa < -1.0f) tmpWeigth[i].sa = -1.0f;
                        //tmpWeigth[i].sa *= maxDeviation;
                        break;
                    case 11:
                        
                        tmpWeigth[i].sb += AdjusRatio; //if(tmpWeigth[i].sb > 1.0f) tmpWeigth[i].sb = 1.0f; else if(tmpWeigth[i].sb < -1.0f) tmpWeigth[i].sb = -1.0f;
                        //tmpWeigth[i].sb *= maxDeviation;
                        break;
                    case 12:
                        
                        tmpWeigth[i].sc += AdjusRatio; //if(tmpWeigth[i].sc > 1.0f) tmpWeigth[i].sc = 1.0f; else if(tmpWeigth[i].sc < -1.0f) tmpWeigth[i].sc = -1.0f;
                        //tmpWeigth[i].sc *= maxDeviation;
                        break;
                    case 13:
                        
                        tmpWeigth[i].sd += AdjusRatio; //if(tmpWeigth[i].sd > 1.0f) tmpWeigth[i].sd = 1.0f; else if(tmpWeigth[i].sd < -1.0f) tmpWeigth[i].sd = -1.0f;
                        //tmpWeigth[i].sd *= maxDeviation;
                        break;
                    case 14:
                        
                        tmpWeigth[i].se += AdjusRatio; //if(tmpWeigth[i].se > 1.0f) tmpWeigth[i].se = 1.0f; else if(tmpWeigth[i].se < -1.0f) tmpWeigth[i].se = -1.0f;
                        //tmpWeigth[i].se *= maxDeviation;
                        break;
                    case 15:
                        
                        tmpWeigth[i].sf += AdjusRatio; //if(tmpWeigth[i].sf > 1.0f) tmpWeigth[i].sf = 1.0f; else if(tmpWeigth[i].sf < -1.0f) tmpWeigth[i].sf = -1.0f;
                        //tmpWeigth[i].sf *= maxDeviation;
                        break;
                }
                
                //barrier(CLK_LOCAL_MEM_FENCE);
                /*if(isnan(tmpWeigth[ i ].s0)){
                    a.s0 += tmpWeigth[ i ].s0 * InNodeVar.s0;    //输入节点数据与权重相乘
                } if(isnan(tmpWeigth[ i ].s1)){
                    a.s1 += tmpWeigth[ i ].s1 * InNodeVar.s1;    //输入节点数据与权重相乘
                } if(isnan(tmpWeigth[ i ].s2)){
                    a.s2 += tmpWeigth[ i ].s2 * InNodeVar.s2;    //输入节点数据与权重相乘
                } if(isnan(tmpWeigth[ i ].s3)){
                    a.s3 += tmpWeigth[ i ].s3 * InNodeVar.s3;    //输入节点数据与权重相乘
                } if(isnan(tmpWeigth[ i ].s4)){
                    a.s4 += tmpWeigth[ i ].s4 * InNodeVar.s4;    //输入节点数据与权重相乘
                } if(isnan(tmpWeigth[ i ].s5)){
                    a.s5 += tmpWeigth[ i ].s5 * InNodeVar.s5;    //输入节点数据与权重相乘
                } if(isnan(tmpWeigth[ i ].s6)){
                    a.s6 += tmpWeigth[ i ].s6 * InNodeVar.s6;    //输入节点数据与权重相乘
                } if(isnan(tmpWeigth[ i ].s7)){
                    a.s7 += tmpWeigth[ i ].s7 * InNodeVar.s7;    //输入节点数据与权重相乘
                } if(isnan(tmpWeigth[ i ].s8)){
                    a.s8 += tmpWeigth[ i ].s8 * InNodeVar.s8;    //输入节点数据与权重相乘
                } if(isnan(tmpWeigth[ i ].s9)){
                    a.s9 += tmpWeigth[ i ].s9 * InNodeVar.s9;    //输入节点数据与权重相乘
                } if(isnan(tmpWeigth[ i ].sa)){
                    a.sa += tmpWeigth[ i ].sa * InNodeVar.sa;    //输入节点数据与权重相乘
                } if(isnan(tmpWeigth[ i ].sb)){
                    a.sb += tmpWeigth[ i ].sb * InNodeVar.sb;    //输入节点数据与权重相乘
                } if(isnan(tmpWeigth[ i ].sc)){
                    a.sc += tmpWeigth[ i ].sc * InNodeVar.sc;    //输入节点数据与权重相乘
                } if(isnan(tmpWeigth[ i ].sd)){
                    a.sd += tmpWeigth[ i ].sd * InNodeVar.sd;    //输入节点数据与权重相乘
                } if(isnan(tmpWeigth[ i ].se)){
                    a.se += tmpWeigth[ i ].se * InNodeVar.se;    //输入节点数据与权重相乘
                } if(isnan(tmpWeigth[ i ].sf)){
                    a.sf += tmpWeigth[ i ].sf * InNodeVar.sf;    //输入节点数据与权重相乘
                }*/
                
                a += tmpWeigth[ i ] * InNodeVar;    //输入节点数据与权重相乘
                //printf("tmpWeigth %#v16f \n", tmpWeigth[ i ]);
            }
            ++modify_wid;



            if(isnan(a.s0) != 0){
                a.s0 = -1.0;
            } if(isnan(a.s1) != 0){
                a.s1 = -1.0;    //输入节点数据与权重相乘
            } if(isnan(a.s2) != 0){
                a.s2 = -1.0;    //输入节点数据与权重相乘
            } if(isnan(a.s3) != 0){
                a.s3 = -1.0;    //输入节点数据与权重相乘
            } if(isnan(a.s4) != 0){
                a.s4 = -1.0;    //输入节点数据与权重相乘
            } if(isnan(a.s5) != 0){
                a.s5 = -1.0;    //输入节点数据与权重相乘
            } if(isnan(a.s6) != 0){
                a.s6 = -1.0;    //输入节点数据与权重相乘
            } if(isnan(a.s7) != 0){
                a.s7 = -1.0;    //输入节点数据与权重相乘
            } if(isnan(a.s8) != 0){
                a.s8 = -1.0;    //输入节点数据与权重相乘
            } if(isnan(a.s9) != 0){
                a.s9 = -1.0;    //输入节点数据与权重相乘
            } if(isnan(a.sa) != 0){
                a.sa = -1.0;    //输入节点数据与权重相乘
            } if(isnan(a.sb) != 0){
                a.sb = -1.0;    //输入节点数据与权重相乘
            } if(isnan(a.sc) != 0){
                a.sc = -1.0;    //输入节点数据与权重相乘
            } if(isnan(a.sd) != 0){
                a.sd = -1.0;    //输入节点数据与权重相乘
            } if(isnan(a.se) != 0){
                a.se = -1.0;    //输入节点数据与权重相乘
            } if(isnan(a.sf) != 0){
                a.sf = -1.0;    //输入节点数据与权重相乘
            }
            
            
            //printf("a %#v16f \n", a);
            //printf("InNodeVar %#v16f \n", InNodeVar);
            float Exp = (-(a.s0+a.s1+a.s2+a.s3+a.s4+a.s5+a.s6+a.s7+a.s8+a.s9+a.sa+a.sb+a.sc+a.sd+a.se+a.sf+-1));
            //if(isinf(Exp)) Exp = 100.0;
            //if(isnan(new_outNodeVar)) new_outNodeVar = -1.0;
            float new_outNodeVar = 0.0;
            new_outNodeVar = 1.0f / (1.0f + exp(Exp) * (1.0f-maxDeviation));
            //new_outNodeVar *= 1.0f-maxDeviation;
            //new_outNodeVar = 1.0f / (1.0f + Exp);
            //new_outNodeVar = (1.0f-exp(Exp)) / (1.0f + exp(Exp));
            
            
            //printf("Exp %f \n", Exp);
            //printf("new_outNodeVar %f \n", new_outNodeVar);
            //printf("outNodeVar %f \n", outNodeVar);

            float x = fabs(new_outNodeVar) - fabs(outNodeVar);
            
            //if(layer) for(uint i=0; i<16; ++i) printf("nodeVarA %f \n", nodeVarA[ i ]);
            if(x < 0.000000000000001f ){
                break;
                return;
            }
            
            
            //printf("outNodeVar %i \n", isnan(0.0001));
            AdjusRatio = (outNodeVar - new_outNodeVar) * LearningRate;
            //AdjusRatio = sign(outNodeVar - new_outNodeVar) * LearningRate * (1.0-maxDeviation);
            //AdjusRatio = -LearningRate * (1.0-maxDeviation);
            //AdjusRatio = LearningRate * (maxDeviation - 1.0);
            //AdjusRatio = (1.0-maxDeviation);
            //AdjusRatio = (outNodeVar / new_outNodeVar);
            
            //printf("maxDeviation %f \n",  AdjusRatio );
            //printf("AdjusRatio %f \n", maxDeviation);
            //printf("AdjusRatio %f \n", exp(1.0));
            //printf("new_outNodeVar %#v4f \n", (float4)(outNodeVar, new_outNodeVar, outNodeVar - new_outNodeVar, AdjusRatio));
        }
        
        WeightModifyID[WeightOffset_size[layer].s1 + groupID] = modify_wid;
        

        //------------------------------跳到上一层--------------------------------------
        //if(layer && groupID == 0){ }
        
    }
    
    //for(i=0; i<BlockNum; ++i) WeightsNode[weightsID] = tmpWeigth[i];

    barrier(CLK_LOCAL_MEM_FENCE);
    WeightsNode[weightsID] = tmpWeigth[localID];
    //return;
    if(layer && groupID == 0 && localID == 0){

        ushort GroupNum_OutLayer = WeightOffset_size[layer].s1;
        ushort LocalBlockNum = WeightOffset_size[layer-1].s1/16;
        if(!LocalBlockNum) LocalBlockNum = 1;
        //for(uint i=0; i<16; ++i) printf("AdjusRatio %f \n", nodeVarA[ i ]);

        ndrange_t ndrange1 = ndrange_1D(GroupNum_OutLayer);
        //printf("maxDeviation %f \n", maxDeviation);
        void(^func_Block)(void) = ^{ f_aiHiddenLayerFeedforward(nodeVarA, nodeVarB, WeightsNode, WeightOffset_size, nodeVar_Buf, WeightModifyID, maxDeviation, LearningRate, LocalBlockNum, layer-1); } ;
        int err_ret = enqueue_kernel(get_default_queue(), CLK_ENQUEUE_FLAGS_WAIT_KERNEL, ndrange1, func_Block);

    }
    
    

}




